#ifndef	linear_included
#define	linear_included

template<uint32_t	I,	uint32_t	O,	uint32_t	C>
struct	linear {
    parameter<I*O>	wei;

    linear() {
        wei.rand(_C2F);
    }

    void	save(FILE	*F) {
        wei.save(F);
    }

    void	load(FILE	*F) {
        wei.load(F);
    }

    uint32_t	size(void) {
        return	wei.size();
    }

    void	fw(tensor<half,I*C>	&inp,	tensor<half,O*C>	&out,	float	beta) {
        half	alf=__float2half_rn(1/sqrtf(I)),bet=__float2half_rn(beta);
        wei.z2w();
        cublasGemmEx(handle,CUBLAS_OP_T,CUBLAS_OP_N,O,C,I,&alf,wei.w.p,CUDA_R_16F,I,inp.p,CUDA_R_16F,I,&bet,out.p,CUDA_R_16F,O,CUBLAS_COMPUTE_16F,CUBLAS_GEMM_DEFAULT);
    }

    void	fwt(tensor<half,I*C>	&inp,	tensor<half,O*C>	&out,	float	beta) {
        half	alf=__float2half_rn(1/sqrtf(I)),bet=__float2half_rn(beta);
        wei.z2w();
        cublasGemmEx(handle,CUBLAS_OP_T,CUBLAS_OP_N,C,O,I,&alf,inp.p,CUDA_R_16F,I,wei.w.p,CUDA_R_16F,I,&bet,out.p,CUDA_R_16F,C,CUBLAS_COMPUTE_16F,CUBLAS_GEMM_DEFAULT);
    }

    void	bk(void) {
        wei.m2z();
    }
};

#endif
